#include <stdio.h>
#include <cuda.h>

#define GB	(1024*1024*1024)


#define CUDA_SAFE_CALL_NO_SYNC( call) do {                                \
    cudaError err = call;                                                    \
    if( cudaSuccess != err) {                                                \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",        \
                __FILE__, __LINE__, cudaGetErrorString( err) );              \
        exit(EXIT_FAILURE);                                                  \
    } } while (0)

#define CUDA_SAFE_CALL( call) do {                                        \
    CUDA_SAFE_CALL_NO_SYNC(call);                                            \
    cudaError err = cudaThreadSynchronize();                                 \
    if( cudaSuccess != err) {                                                \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",        \
                __FILE__, __LINE__, cudaGetErrorString( err) );              \
        exit(EXIT_FAILURE);                                                  \
    } } while (0)

__global__ static void measure_seq_read_int(char * buf, int tupleNum ,char *out){
	int startIndex = blockIdx.x*blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for(int i=startIndex;i<tupleNum;i+=stride){
		((int*)out)[i] = ((int*)buf)[i];
	}
} 

__global__ static void measure_seq_write_int(char * buf, int tupleNum){
	int startIndex = blockIdx.x*blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;

	for(int i=startIndex;i<tupleNum;i+=stride){
		((int *)buf)[i] = i;
	}
}


__global__ static void measure_write_char(int tupleNum, int attrSize, char *outBuf){
	int startIndex = blockIdx.x*blockDim.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;
	const char *buf = "abcdefghijklmnopqrstuvwxyz";

	for(int i=startIndex;i<tupleNum;i+=stride){
		memcpy(outBuf + i*attrSize, buf , attrSize);
	}
}

int main( int argc, char **argv){
	int res = 0;
	float bw;

	char * testReadMem;
	char * testWriteMem;

	dim3 grid(1024);
	dim3 block(1024);

	float gpuTime;
	cudaEvent_t startGPU, stopGPU;
        cudaEventCreate(&startGPU);
        cudaEventCreate(&stopGPU);

// test int first
//test sequential read and write first

	int tupleNum = 1024*1024 * 64; 
	char * cpuMem = (char *) malloc(tupleNum*sizeof(int)); 
	if(!cpuMem){
		printf("Failed to allocate host memory\n");
		exit(-1);
	}

	for(int i=0;i<tupleNum;i++){
		((int *)cpuMem)[i] = i;	
	}

	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&testReadMem, sizeof(int) * tupleNum));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&testWriteMem, sizeof(int) * tupleNum));
	CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(testReadMem, cpuMem, sizeof(int)* tupleNum, cudaMemcpyHostToDevice));

	cudaEventRecord(startGPU,0);
	measure_seq_write_int<<<grid,block>>>(testWriteMem,tupleNum);
	CUDA_SAFE_CALL_NO_SYNC(cudaDeviceSynchronize());
	cudaEventRecord(stopGPU,0);
        cudaEventSynchronize(stopGPU);
        cudaEventElapsedTime(&gpuTime,startGPU,stopGPU);

	CUDA_SAFE_CALL_NO_SYNC(cudaFree(testReadMem));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(testWriteMem));

	bw = tupleNum * sizeof(int)/ (gpuTime * GB / 1000); 
	printf("sequential read int bandwidth(GB/s):%f\n",bw);


	for(int i=1;i<=16;i++){
		CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&testWriteMem, i * tupleNum));

		cudaEventRecord(startGPU,0);
		measure_write_char<<<grid,block>>>(tupleNum, i, testWriteMem);
		CUDA_SAFE_CALL_NO_SYNC(cudaDeviceSynchronize());
		cudaEventRecord(stopGPU,0);
        	cudaEventSynchronize(stopGPU);
        	cudaEventElapsedTime(&gpuTime,startGPU,stopGPU);

		CUDA_SAFE_CALL_NO_SYNC(cudaFree(testWriteMem));
		bw = tupleNum * i/ (gpuTime * GB / 1000); 
		printf("sequential write %d char bandwidth(GB/s):%f\n", i, bw);
	}

	return res;
}
